/*
 *  Copyright 2008-2010 NVIDIA Corporation
 *
 *  Licensed under the Apache License, Version 2.0 (the "License");
 *  you may not use this file except in compliance with the License.
 *  You may obtain a copy of the License at
 *
 *      http://www.apache.org/licenses/LICENSE-2.0
 *
 *  Unless required by applicable law or agreed to in writing, software
 *  distributed under the License is distributed on an "AS IS" BASIS,
 *  WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied.
 *  See the License for the specific language governing permissions and
 *  limitations under the License.
 */


/*! \file dereference.h
 *  \brief Overloads for everything dereferenceable by dereference.h
 */

#pragma once

#include <thrust/detail/type_traits.h>

namespace thrust {

// forward declarations of iterator types
template<typename T>
class device_ptr;

namespace detail {

template<typename Pointer>
class normal_iterator;

template<typename Iterator, typename Space>
class forced_iterator;

} // end detail


template < typename Value,
         typename Incrementable,
         typename Space >
class constant_iterator;


template <class UnaryFunc, class Iterator, class Reference, class Value>
class transform_iterator;

template < typename Incrementable,
         typename Space,
         typename Traversal,
         typename Difference >
class counting_iterator;

template <typename IteratorTuple>
class zip_iterator;

template <typename BidirectionalIterator>
class reverse_iterator;

template <typename ElementIterator, typename IndexIterator>
class permutation_iterator;

namespace detail {

namespace device {

template <typename> struct dereference_result {};

// specialize dereference_result for T*
template<typename T>
struct dereference_result<T*> {
    typedef T& type;
}; // end dereference_result

// raw pointer
template<typename T>
inline __host__ __device__
typename dereference_result<T*>::type
dereference(T* ptr) {
    return *ptr;
} // dereference

template<typename T, typename IndexType>
inline __host__ __device__
typename dereference_result<T*>::type
dereference(T* ptr, IndexType n) {
    return ptr[n];
} // dereference



// device_ptr prototypes
template<typename T>
inline __host__ __device__
typename dereference_result< device_ptr<T> >::type
dereference(device_ptr<T> ptr);

template<typename T, typename IndexType>
inline __host__ __device__
typename dereference_result< device_ptr<T> >::type
dereference(thrust::device_ptr<T> ptr, IndexType n);



// normal_iterator
template<typename T>
inline __host__ __device__
typename dereference_result< normal_iterator< device_ptr<T> > >::type
dereference(const normal_iterator< device_ptr<T> >& iter);

template<typename T, typename IndexType>
inline __host__ __device__
typename dereference_result< normal_iterator< device_ptr<T> > >::type
dereference(const normal_iterator< device_ptr<T> >& iter, IndexType n);


// forced_iterator
template<typename Iterator, typename Space>
inline __host__ __device__
typename dereference_result< forced_iterator< Iterator, Space > >::type
dereference(const forced_iterator< Iterator, Space >& iter);

template<typename Iterator, typename Space, typename IndexType>
inline __host__ __device__
typename dereference_result< forced_iterator< Iterator, Space > >::type
dereference(const forced_iterator< Iterator, Space >& iter, IndexType n);



// transform_iterator prototypes
template<typename UnaryFunc, typename Iterator, typename Reference, typename Value>
inline __host__ __device__
typename dereference_result< thrust::transform_iterator<UnaryFunc, Iterator, Reference, Value> >::type
dereference(const thrust::transform_iterator<UnaryFunc, Iterator, Reference, Value>& iter);

template<typename UnaryFunc, typename Iterator, typename Reference, typename Value, typename IndexType>
inline __host__ __device__
typename dereference_result< thrust::transform_iterator<UnaryFunc, Iterator, Reference, Value> >::type
dereference(const thrust::transform_iterator<UnaryFunc, Iterator, Reference, Value>& iter, IndexType n);



// counting_iterator prototypes
template<typename Incrementable, typename Space, typename Traversal, typename Difference>
inline __host__ __device__
typename dereference_result< thrust::counting_iterator<Incrementable, Space, Traversal, Difference> >::type
dereference(const thrust::counting_iterator<Incrementable, Space, Traversal, Difference>& iter);

template<typename Incrementable, typename Space, typename Traversal, typename Difference, typename IndexType>
inline __host__ __device__
typename dereference_result< thrust::counting_iterator<Incrementable, Space, Traversal, Difference> >::type
dereference(const thrust::counting_iterator<Incrementable, Space, Traversal, Difference>& iter, IndexType n);


// constant_iterator prototypes
template<typename Value, typename Incrementable, typename Space>
inline __host__ __device__
typename dereference_result< thrust::constant_iterator<Value, Incrementable, Space> >::type
dereference(const thrust::constant_iterator<Value, Incrementable, Space>& iter);

template<typename Value, typename Incrementable, typename Space, typename IndexType>
inline __host__ __device__
typename dereference_result< thrust::constant_iterator<Value, Incrementable, Space> >::type
dereference(const thrust::constant_iterator<Value, Incrementable, Space>& iter, IndexType n);


// zip_iterator prototypes
template<typename IteratorTuple>
inline __host__ __device__
typename dereference_result< thrust::zip_iterator<IteratorTuple> >::type
dereference(const thrust::zip_iterator<IteratorTuple>& iter);

template<typename IteratorTuple, typename IndexType>
inline __host__ __device__
typename dereference_result< thrust::zip_iterator<IteratorTuple> >::type
dereference(const thrust::zip_iterator<IteratorTuple>& iter, IndexType n);


// reverse_iterator prototypes
template<typename BidirectionalIterator>
inline __host__ __device__
typename dereference_result< thrust::reverse_iterator<BidirectionalIterator> >::type
dereference(const thrust::reverse_iterator<BidirectionalIterator>& iter);

template<typename BidirectionalIterator, typename IndexType>
inline __host__ __device__
typename dereference_result< thrust::reverse_iterator<BidirectionalIterator> >::type
dereference(const thrust::reverse_iterator<BidirectionalIterator>& iter, IndexType n);


// permutation_iterator prototypes
template<typename ElementIterator, typename IndexIterator>
inline __host__ __device__
typename dereference_result< thrust::permutation_iterator<ElementIterator, IndexIterator> >::type
dereference(const thrust::permutation_iterator<ElementIterator, IndexIterator>& iter);

template<typename ElementIterator, typename IndexIterator, typename IndexType>
inline __host__ __device__
typename dereference_result< thrust::permutation_iterator<ElementIterator, IndexIterator> >::type
dereference(const thrust::permutation_iterator<ElementIterator, IndexIterator>& iter, IndexType n);

} // end device

} // end detail

} // end thrust

